#include <cuda_runtime.h>
#include <iostream>

__global__ void vec_inc(float *data, int size) {
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	if (idx < size) {
		data[idx] +=1; // 将数组元素值+1
	}
}

int main() {
	// 创建CUDA流，指定非阻塞标志
	cudaStream_t stream;
	cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

	// 分配和初始化主机内存
	int num = 10;
	size_t size = num * sizeof(float);
	float *hostArray = new float[num];
	for (int i = 0; i < num; i++) {
		hostArray[i] = float(i);
	}

	// 分配设备内存
	float *deviceArray;
	cudaMalloc((void **)&deviceArray, size);

	// 异步从主机到设备的内存复制
	cudaMemcpyAsync(deviceArray, hostArray, size, cudaMemcpyHostToDevice, stream);

	// 定义核函数的执行配置
	int blockSize = 256;
	int numBlocks = (num + blockSize - 1) / blockSize;

	// 在流上异步执行核函数
	vec_inc<<<numBlocks, blockSize, 0, stream>>>(deviceArray, num);

	// 等待流中的所有操作完成
	cudaStreamSynchronize(stream);

	// 将结果从设备异步复制回主机
	cudaMemcpyAsync(hostArray, deviceArray, size, cudaMemcpyDeviceToHost, stream);
	cudaStreamSynchronize(stream);

	// 检查结果（如果需要）
	for (int i = 0; i < num; i++) {
		printf("%f ", hostArray[i]);
	}

	// 清理
	cudaStreamDestroy(stream);
	cudaFree(deviceArray);
	delete[] hostArray;

	std::cout << "Done" << std::endl;
	return 0;
}
